/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/config.h>

namespace cuqu
{
    namespace detail
    {
        struct barrier_lockfree {
            enum { 
                grid_size_max = 1024,
                array_size = grid_size_max
            };
            int last_goal;
            int arrayIn[array_size];
            int arrayOut[array_size];           
        };

        //typedef struct barrier_lockfree barrier_t;

#define CUQU_INIT_BARRIER_LOCKFREE { 0, {0,}, {0,} }

        static __inline__ __device__ int barrier_wait(barrier_lockfree *gbl_barr, int n_waiters)
        {
            const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
            const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
            const uint block_size = blockDim.x*blockDim.y*blockDim.z;
            const uint gridSize = gridDim.y*gridDim.x;
            const uint gid = tid + bid*block_size;
            int cnt;
            int *arrayIn = gbl_barr->arrayIn;
            int *arrayOut = gbl_barr->arrayOut;
            int *last_goal = &gbl_barr->last_goal;
            int goal;
            int ret = SUCCESS;
            const int retry_max = 500000;

#if defined(CUQU_DEBUG)
            if(gridSize > barrier_lockfree::array_size) {
                cuqu_error("invalid array_size!!\n");
                return INVALID;
            }
#endif

            //__threadfence();
            //goal = n_waiters;
            goal = *(volatile int*)last_goal + n_waiters;

            // all blocks signal toward block 1
            if(0 == tid) {
                arrayIn[bid] = goal;
                __threadfence();
            }
            //__threadfence();
            //__syncthreads();

            // block 0 waits for signals toward itself
            // and signals back
            if(0 == bid) {
                if(tid < gridSize) {
                    int retry = retry_max;
                    do {
                        int in = *(volatile int*)&arrayIn[tid];
                        if(in == goal)
                            break;
                        if(--retry < 0) {
                            cuqu_error("[%d:%d] timeout polling arrayIn[%d]=%d for goal=%d\n", bid, tid, tid, in, goal);
                            ret = WOULDBLOCK;
                            break;
                        }
                    } while(true);
                }
                // update last_goal
                if(0 == tid) {
                    *last_goal = goal;
                    //__threadfence();
                    //cuqu_error("%d:%d goal=%d sync done\n", bid, tid, goal);
                }
                // bid 0 threads are synch'd
                __syncthreads();
                // bid 0 sends back a signal towards all other blocks
                if(tid < gridSize) {
                    *(volatile int*)&arrayOut[tid] = goal;
                    __threadfence();
                }
           }
            // all blocks wait for signal from block 0
            if(0 == tid) {
                int retry = retry_max;
                do {
                    int out = *(volatile int*)&arrayOut[bid];
                    if(out == goal)
                        break;
                    if(--retry < 0) {
                        cuqu_error("[%d:%d] timeout polling arrayOut[%d]=%d for goal=%d\n", bid, tid, bid, out, goal);
                        ret = WOULDBLOCK;
                        break;
                    }
                } while(true);
            }
            __syncthreads();
            return ret;
        }

    } // end namespace detail

} // end namespace cuqu

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
